#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

using namespace std;


// constant memory位于global memory，但是为只读的。作用范围是：所有线程+host
// 直接定义和初始化（可选）
__constant__ float warp_matrix[6] = {1, 2, 3, 4, 5, 6};

static __global__ void testKernel() {
    // 在核函数内，constant memory不能修改，否则报错
    printf("warp_matrix[%d] = %f\n", threadIdx.x, warp_matrix[threadIdx.x]);
}

int main() {
    // 覆盖warp_matrix的初始化值定义
    float host_warp_matrix[6] = {6, 5, 4, 3, 2, 1};
    // cudaMemcpyToSymbol 是 CUDA（Compute Unified Device Architecture，统一计算设备架构）编程中的一个重要函数，
    // 主要用于将主机（CPU）内存中的数据复制到设备（GPU）的全局符号（如全局变量、常量等）对应的内存空间。
    // warp_matrix是显存global memory中的常量区域，可以用cudaMemcpyToSymbol执行h2d写操作，但是在核函数中不能写。
    cudaMemcpyToSymbol(warp_matrix, host_warp_matrix, sizeof(host_warp_matrix));
    testKernel<<<1, 6>>>();
    cudaDeviceSynchronize();
    return 0;
}


// warp_matrix[0] = 6.000000
// warp_matrix[1] = 5.000000
// warp_matrix[2] = 4.000000
// warp_matrix[3] = 3.000000
// warp_matrix[4] = 2.000000
// warp_matrix[5] = 1.000000
